#ifndef	simd8_included
#define	simd8_included

#define	_F2C	16.0f
#define	_C2F	0.0625f

struct	half8 {
    __nv_half2	x,y,z,w;
};

struct	float8 {
    float2	x,y,z,w;
};

struct	char8 {
    char2	x,y,z,w;
};

__device__ inline	half8	_f2b8(float8	x) {
    half8	y;
    y.x=__float22half2_rn(x.x);
    y.y=__float22half2_rn(x.y);
    y.z=__float22half2_rn(x.z);
    y.w=__float22half2_rn(x.w);
    return	y;
}

__device__ inline	float8	_b2f8(half8	x) {
    float8	y;
    y.x=__half22float2(x.x);
    y.y=__half22float2(x.y);
    y.z=__half22float2(x.z);
    y.w=__half22float2(x.w);
    return	y;
}

__device__ inline	char2	_f2c(float2	x) {
    char2	y;
    y.x=roundf(fminf(fmaxf((x.x*_F2C),-127.0f),127.0f));
    y.y=roundf(fminf(fmaxf((x.y*_F2C),-127.0f),127.0f));
    return	y;
}

__device__ inline	float2	_c2f(char2	x) {
    float2	y;
    y.x=x.x*_C2F;
    y.y=x.y*_C2F;
    return	y;
}

__device__ inline	char8	_f2c8(float8	x) {
    char8	y;
    y.x=_f2c(x.x);
    y.y=_f2c(x.y);
    y.z=_f2c(x.z);
    y.w=_f2c(x.w);
    return	y;
}

__device__ inline	float8	_c2f8(char8	x) {
    float8	y;
    y.x=_c2f(x.x);
    y.y=_c2f(x.y);
    y.z=_c2f(x.z);
    y.w=_c2f(x.w);
    return	y;
}

__device__ inline	float	_fdot8(float8	x,	float8	y) {
    return	x.x.x*y.x.x+x.x.y*y.x.y+x.y.x*y.y.x+x.y.y*y.y.y+x.z.x*y.z.x+x.z.y*y.z.y+x.w.x*y.w.x+x.w.y*y.w.y;
}

__device__ inline	float	_fsum8(float8	x) {
    return	x.x.x+x.x.y+x.y.x+x.y.y+x.z.x+x.z.y+x.w.x+x.w.y;
}

__device__ inline	float8	_fmul18(float8	x,	float	y) {
    x.x.x*=y;
    x.x.y*=y;
    x.y.x*=y;
    x.y.y*=y;
    x.z.x*=y;
    x.z.y*=y;
    x.w.x*=y;
    x.w.y*=y;
    return	x;
}

__device__ inline	float8	_fadd18(float8	x,	float	y) {
    x.x.x+=y;
    x.x.y+=y;
    x.y.x+=y;
    x.y.y+=y;
    x.z.x+=y;
    x.z.y+=y;
    x.w.x+=y;
    x.w.y+=y;
    return	x;
}

__device__ inline	float8	_fadd8(float8	x,	float8	y) {
    float8	z;
    z.x.x=x.x.x+y.x.x;
    z.x.y=x.x.y+y.x.y;
    z.y.x=x.y.x+y.y.x;
    z.y.y=x.y.y+y.y.y;
    z.z.x=x.z.x+y.z.x;
    z.z.y=x.z.y+y.z.y;
    z.w.x=x.w.x+y.w.x;
    z.w.y=x.w.y+y.w.y;
    return	z;
}

__device__ inline	half8	_hadd8(half8	x,	half8	y) {
    half8	z;
    z.x=__hadd2(x.x,y.x);
    z.y=__hadd2(x.y,y.y);
    z.z=__hadd2(x.z,y.z);
    z.w=__hadd2(x.w,y.w);
    return	z;
}
__device__ inline	half8	_hmul8(half8	x,	half8	y) {
    half8	z;
    z.x=__hmul2(x.x,y.x);
    z.y=__hmul2(x.y,y.y);
    z.z=__hmul2(x.z,y.z);
    z.w=__hmul2(x.w,y.w);
    return	z;
}

__device__ inline	half8	_hmul18(half8	x,	half	m) {
    half8	z;
    half2	y= {m,m};
    z.x=__hmul2(x.x,y);
    z.y=__hmul2(x.y,y);
    z.z=__hmul2(x.z,y);
    z.w=__hmul2(x.w,y);
    return	z;
}
__device__ inline	half8	_hexp8(half8	x) {
    half8	z;
    z.x=h2exp(x.x);
    z.y=h2exp(x.y);
    z.z=h2exp(x.z);
    z.w=h2exp(x.w);
    return	z;
}
#endif
